/*
 *  Copyright 2010 INFN - APE group
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#pragma once

#include <cuqu/config.h>

namespace cuqu
{
    namespace detail
    {
        struct barrier_oneloop {
            atomic_t m_cnt;
        };
        
        //typedef struct barrier barrier_t;

#define CUQU_INIT_BARRIER_ONELOOP { 0 }

        static __inline__ __device__ int barrier_wait(barrier_oneloop *gbl_barrier, int n_waiters)
        {
            const unsigned int tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
            const unsigned int bid = blockIdx.x + gridDim.x*blockIdx.y;
            //const unsigned int block_size = blockDim.x*blockDim.y*blockDim.z;
            //const unsigned int grid_size = gridDim.y*gridDim.x;
            //const unsigned int gid = tid + bid*block_size;
            int cnt;
            int retry = 200000;
            int retcode = SUCCESS;

            // this fence adds overhead!!
            // from 2.7us to 1.6us per invocation
            //__threadfence_system();

            // every block does mod_inc
            if(0 == tid) {
                //cuqu_trace("[%d:%d] barrier %p  mod_inc()\n", bid, tid, gbl_barrier);
                unsigned int oldcnt = atomic_mod_inc(&gbl_barrier->m_cnt, n_waiters-1);

#if defined(CUQU_DEBUG)
                if(oldcnt >= n_waiters)
                    cuqu_error("ERROR invalid barrier counter %d\n", oldcnt);
#endif

                // every block waits till the counter reaches 0 due to mod_inc
                do {
                    cnt = atomic_read(&gbl_barrier->m_cnt);
                    // != n_waiters, due to mod_inc
                    if(0 == cnt) {
                        //cuqu_trace("[%d:%d] barrier %p done, counter reached %d\n", bid, tid, gbl_barrier, cnt);
                        break;
                    }
#if defined(CUQU_DEBUG)
                    else if(cnt >= n_waiters) {
                        cuqu_error("[%d:%d] ERROR invalid barrier counter %d\n", bid, tid, cnt);
                        retcode = INVALID;
                        break;
                    }
#endif
                    --retry;
                    if(retry < 0) {
                        cuqu_error("[%d:%d] ERROR barrier %p timeout polling counter, cnt=%d n_waiters=%d\n", bid, tid, gbl_barrier, cnt, n_waiters);
                        retcode = WOULDBLOCK;
                        break;
                    }
                } while(true);
            }
            __syncthreads();

            return retcode;
        }

    } // end namespace detail
    
} // end namespace cuqu

/*
 * Local variables:
 *  mode: c++
 *  c-indent-level: 4
 *  c-basic-offset: 4
 *  tab-width: 4
 *  indent-tabs-mode: nil
 * End:
 */
